-
Notifications
You must be signed in to change notification settings - Fork 3
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
BitwiseAND and BitwiseOR on HOST and HIP #230
BitwiseAND and BitwiseOR on HOST and HIP #230
Conversation
Also includes fixing f16 and f32 datatype of HOST
Update Copywrite
Includes reference output
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
@snehaa8 Pls address comments. Replicate changes mentioned in either bitwiseOr/AND for both.
Lets combine BitwiseAND + BitwiseOR in this Internal PR and close the other one for ease. You can work on this sn/bitwise_OR branch for both.
@@ -0,0 +1,957 @@ | |||
/* | |||
MIT License |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Replace exact text with the blank lines as in LICENSE file outside:
MIT License
Copyright (c) 2019 - 2024 Advanced Micro Devices, Inc.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in all
copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
SOFTWARE.
{ | ||
__m128i pxSrc[8]; | ||
__m128i pxMask = _mm_setr_epi8(0, 3, 6, 9, 1, 4, 7, 10, 2, 5, 8, 11, 12, 13, 14, 15); | ||
__m128i pxMaskRGB = _mm_setr_epi8(0, 4, 8, 12, 2, 6, 10, 14, 1, 5, 9, 13, 3, 7, 11, 15); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Don't we have the 0,3,6,9 mask or the 0,4,8,12 mask pre-allocated outside of the runtime execution path somewhere since they are common?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Checked this while implementing, didn't find any in the fashion I needed.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Should i add this at the start of rpp_cpu_simd where other common constants are defined?
#endif | ||
for (; vectorLoopCount < bufferLength; vectorLoopCount += 3) | ||
{ | ||
*dstPtrTempR++ = RPPPIXELCHECKF32((float)((uint)(srcPtr1Temp[0] * 255) & (uint)(srcPtr2Temp[0] * 255)) / 255); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Please add a comment like below on HOST/HIP for bitwiseAND and bitwiseOR to clarify to the reader. Either the link you pointed to before, or establish validity with openCV or other lib.
// BitwiseAND / BitwiseOR are logical operations only on U8/I8 types. For a float / half precision image (pixel values from 0-1), the BitwiseAND / BitwiseOR is applied on a 0-255 range-translated approximation, of the original 0-1 decimal-range image
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Added comments and link as pointed to before
@@ -0,0 +1,957 @@ | |||
/* | |||
MIT License |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Same comment
#endif | ||
for (; vectorLoopCount < bufferLength; vectorLoopCount += 3) | ||
{ | ||
*dstPtrTempR++ = RPPPIXELCHECKF32((float)((uint)(srcPtr1Temp[0] * 255) | (uint)(srcPtr2Temp[0] * 255)) / 255); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Same comment for Bitwise OR
#include "rpp_hip_common.hpp" | ||
|
||
template <typename T> | ||
__device__ void bitwise_or_hip_compute(T *srcPtr, d_float8 *src1_f8, d_float8 *src2_f8, d_float8 *dst_f8) |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
template <typename T>
__device__ void bitwise_or_hip_compute(T *srcPtr, d_float8 *src1_f8, d_float8 *src2_f8, d_float8 *dst_f8)
{
if constexpr ((std::is_same<T, float>::value) || (std::is_same<T, half>::value))
{
rpp_hip_math_multiply8_const(src1_f8, src1_f8, (float4)255);
rpp_hip_math_multiply8_const(src2_f8, src2_f8, (float4)255);
rpp_hip_math_bitwiseOr8(src1_f8, src2_f8, dst_f8);
rpp_hip_math_multiply8_const(dst_f8, dst_f8, (float4)ONE_OVER_255);
}
else if constexpr (std::is_same<T, signed char>::value)
{
rpp_hip_math_add8_const(src1_f8, src1_f8, (float4)128);
rpp_hip_math_add8_const(src2_f8, src2_f8, (float4)128);
rpp_hip_math_bitwiseOr8(src1_f8, src2_f8, dst_f8);
rpp_hip_math_subtract8_const(dst_f8, dst_f8, (float4)128);
}
}
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Modified it a little more like
template <typename T>
__device__ void bitwise_and_hip_compute(T *srcPtr, d_float8 *src1_f8, d_float8 *src2_f8, d_float8 *dst_f8)
{
if constexpr ((std::is_same<T, float>::value) || (std::is_same<T, half>::value))
{
rpp_hip_math_multiply8_const(src1_f8, src1_f8, (float4)255);
rpp_hip_math_multiply8_const(src2_f8, src2_f8, (float4)255);
rpp_hip_math_bitwiseAnd8(src1_f8, src2_f8, dst_f8);
rpp_hip_math_multiply8_const(dst_f8, dst_f8, (float4)ONE_OVER_255);
}
else if constexpr (std::is_same<T, signed char>::value)
{
rpp_hip_math_add8_const(src1_f8, src1_f8, (float4)128);
rpp_hip_math_add8_const(src2_f8, src2_f8, (float4)128);
rpp_hip_math_bitwiseAnd8(src1_f8, src2_f8, dst_f8);
rpp_hip_math_subtract8_const(dst_f8, dst_f8, (float4)128);
}
else
rpp_hip_math_bitwiseAnd8(src1_f8, src2_f8, dst_f8);
}
template <typename T> | ||
__device__ void bitwise_or_hip_compute(T *srcPtr, d_float8 *src1_f8, d_float8 *src2_f8, d_float8 *dst_f8) | ||
{ | ||
float4 adjustment_f4; |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Remove unused variable
@@ -97,7 +98,8 @@ std::map<int, string> augmentationMap = | |||
{84, "spatter"}, | |||
{85, "swap_channels"}, | |||
{86, "color_to_greyscale"}, | |||
{87, "tensor_sum"} | |||
{87, "tensor_sum"}, | |||
{92, "bitwise_or"} |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Why is the case number for bitwiseOR this far apart from bitwiseAND? Pls check BatchPD case numbers.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Modified testCase of Bitwise OR to 68 to match with Inclusive OR of BatchPD.
Please take a final look |
@snehaa8 Conflicts. Pull upstream develop into your branch. |
Resolved merge conflicts |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
@snehaa8 I made some minor formatting changes on your branch. Please also move the header files as in comment.
#include <hip/hip_runtime.h> | ||
#include "rpp_hip_common.hpp" | ||
|
||
/* BitwiseAND is logical operation only on U8/I8 types. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
We need to create a "logical_operations" header for bitwise ops. Currently everything is under arithmetic.
Test suite grouping/classification, external .h include, and internal .hpp includes need to change.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Done, please recheck.
Confirmed QA test pass too.
No description provided.